#define SCAN_BLOCK_DIM 256
__kernel void kernel_sat_scan(__global float4 *dev_src, 
							  __global float4 * dev_dst, uint len, uint stride)
{
	uint k = get_global_id(0);

	uint idx = get_local_id(0);
	uint offset = len * get_group_id(1);

	__local float4 src[2*SCAN_BLOCK_DIM];

	src[2*idx] = dev_src[k + offset];
	barrier(CLK_LOCAL_MEM_FENCE); // for coalescing

	if (k >= stride)
		src[2*idx+1] = dev_src[k-stride + offset];
	barrier(CLK_LOCAL_MEM_FENCE); // prevent operating on nonexistent data

	if (k < len) {
		dev_dst[k + offset] = (k >= stride) ? src[2*idx+1] + src[2*idx] : src[2*idx];
	}
}
